导航菜单
首页 >  The OpenCL™ C Specification  > The C++ for OpenCL 1.0 and 2021 Programming Language Documentation

The C++ for OpenCL 1.0 and 2021 Programming Language Documentation

3.3.3. Address space inference

This section details what happens if address spaces for types are notprovided in the source code explicitly. Most of the logic for address spaceinference (i.e. default address space) follows rules from OpenCL C 3.0 s6.7.8.

References inherit rules from pointers and therefore refer to genericaddress space objects by default (see References) ifgeneric address space is supported, otherwise they refer to privateaddress space objects.

Class static data members are deduced to __global address space. Note,that if the C++ for OpenCL 2021 implementation does not support programscope variables (i.e. __opencl_c_program_scope_global_variables featurefrom OpenCL C 3.0 s6.1 is unsupported) the address space qualifier must be specifiedexplicitly and it must be the __constant address space.

All non-static member functions take an implicit object parameter thisthat is a pointer type. By default the this pointer parameter is in thegeneric address space if it is supported and in the private address spaceotherwise. All concrete objects passed as an argument to theimplicit this parameter will be converted to this default (generic or private)address space first if such conversion is valid. Therefore, when member functionsare called with objects created in disjoint address spaces from the default one,the compilation must fail. To prevent the failure the address space on implicitobject parameter this must be specified using address space qualifierson member functions (see Member function qualifier). For example, use of member functions withobjects in __constant address space will always require a __constantmember function qualifier as __constant address space is disjoint with anyother.

Member function qualifiers can also be used in case address space conversionsare undesirable for example for performance reasons. For example, a method canbe implemented to exploit memory access coalescing for segments with memory bank.

Address spaces are not deduced for:

non-pointer/non-reference template parameters except for templatespecializations or non-type based template parameters.

non-pointer/non-reference class members except for static data membersthat are deduced to the __global address space for C++ for OpenCL 1.0or C++ for OpenCL 2021 with the__opencl_c_program_scope_global_variables feature.

non-pointer/non-reference type alias declarations.

decltype expressions.

template void foo() { T m; // address space of 'm' will be known at template instantiation time. T * ptr; // 'ptr' points to generic address space object when it is// supported otherwise to __private address space. T & ref = ...; // 'ref' references an object in generic address space when // it is supported otherwise in __private address space.};template struct S { int i; // 'i' has no address space. static int ii; // 'ii' is in global address space if program scope variables // are supported; otherwise this statement is not legal. int * ptr; // 'ptr' points to int in generic address space if it is supported; // otherwise to __private address space. int & ref = ...; // 'ref' references int in generic address space if it is// supported; otherwise in __private address space.};template void bar(){ S s; // 's' is in __private address space.}struct c1 {};using alias_c1 = c1; // 'alias_c1' is 'c1'.using alias_c1_ptr = c1 *; // 'alias_c1_ptr' is a generic address space pointer to// 'c1' when generic address space is supported; otherwise// it points to 'c1' located in __private address space.__kernel void foo(){ __local int i; decltype(i)* ii; // type of 'ii' is '__local int *__private'.}

For the placeholder type specifier auto an address space of the outer type isdeduced as if it would be any other regular type. However if auto is used in areference or pointer type, the address space of a pointee is taken from the typeof the initialization expression. The logic follows rules for const andvolatile qualifiers.

// This example assumes that generic address space is supported.__kernel void foo(){ __local int i; constexpr int c = 1; __constant auto cai = c; // type of 'cai' is '__constant int' (no deduction). auto aii = cai; // type of 'aii' is '__private int' (regular deduction). auto *ptr = &i; // type of 'ptr' is '__local int * __private' // (addr space of a pointer is deduced regularly, // addr space of its pointee is taken from 'i'). auto *&refptr = ptr; // type of 'refptr' is '__local int * generic & __private'// (addr space of a reference and type of referencing object// is deduced regularly,// addr space of a pointee is taken from the pointee of 'ptr').}

相关推荐: